

//--------------------预定义函数---------------------
void f_RedixSort_OffsetA(	__global ulong2* Hash, __global uint2* Index, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uchar step, uint offset);
void f_RedixSort_OffsetB(	__global ulong2* Hash, __global uint2* Index, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uchar step, uint offset);
void f_RedixSort_Acc(		__global ulong2* Hash, __global uint2* Index, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uchar step);


#define CurrentID ((get_group_id(1)*get_num_groups(0)*get_num_groups(2)) + (get_group_id(2)*get_num_groups(0)) + get_group_id(0))
#define BlockNum 16


__kernel void f_DemainIndexRadixSort(__global ulong2* Hash, __global uint2* HashIndex, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uchar Step){
	const size_t id = CurrentID;
	if(id<eNum){
        
		uint HashOffset = id*BlockNum;
		ulong hashTmp[16];
        uint indexTmp[16];
		uchar countTmp[16];
		uchar mask[16];
		uchar index = 0;
		uchar bitOffset = Step*4;
		
		for(uchar i=0; i<16; ++i){
			Hash[HashOffset+i].s1 = 0x1000000000000;
			hashTmp[i] = Hash[HashOffset+i].s0;
			countTmp[i] = 0;
			mask[i] = 0;
            
            if(!Step){
                HashIndex[HashOffset+i] = HashOffset+i;
            }
            indexTmp[i] = HashIndex[HashOffset+i].s0;
		}
		
		for(uchar i=0; i<16; ++i){
			for(uchar ia=0; ia<16; ++ia){
				if(mask[ia] || hashTmp[ia] >= 0x1000000000000) {
					continue;
				}
				
				uchar bitHash = hashTmp[ia]>>bitOffset;
				bitHash<<=4;
				bitHash>>=4;
				//if(id==0) printf("	b = %u \n", bitHash);
				if(bitHash == i){
					Hash[HashOffset+index].s1 = hashTmp[ia];
                    HashIndex[HashOffset+index].s1 = indexTmp[ia];
					++index;
					++mask[ia];
					countTmp[i] += 1;
				}

				if(index >= BlockNum) break;
			}
			
		}
        
		for(uchar i=0; i<16; ++i){
			rangeOffset[HashOffset+i] = countTmp[i];
			rangeNum[HashOffset+i] = countTmp[i];
		}
		
		//barrier(CLK_GLOBAL_MEM_FENCE);
		
		//偏移计算
		if(id==0){
			size_t work_size[3];
			work_size[0] = get_num_groups(0); 
			work_size[1] = get_num_groups(1); 
			work_size[2] = get_num_groups(2);
			ndrange_t ndrange1 = ndrange_3D(work_size);
			void(^func_SamplingBlock)(void) = ^{ f_RedixSort_OffsetA(Hash, HashIndex, rangeOffset, rangeNum, eNum, Step, 1); };
			int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
			//if(err_ret<0) printf("Previous copy D  = %u \n", err_ret);
			
		}

	}
}


void f_RedixSort_OffsetA(__global ulong2* Hash, __global uint2* HashIndex, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uchar Step, uint offset){
	const size_t id = CurrentID;
	size_t CurrentOffset = id+offset;
	//if(CurrentOffset >= 41597) printf("offset copy D  = %u \n", offset);
	if (CurrentOffset<eNum) {
		CurrentOffset *= BlockNum;
		size_t CurrentBeginOffset = id*BlockNum;
		for(uchar i=0; i<16; ++i){
			rangeOffset[CurrentOffset+i].s0 += rangeOffset[CurrentBeginOffset+i].s1;
		}
	}
	//barrier(CLK_GLOBAL_MEM_FENCE);
	//===============偏移拷贝======================
	if(id==0){
		size_t work_size[3];
		if(offset<=eNum){
			work_size[0] = pow(eNum-offset, 1.0f/3.0f)+1; 
			work_size[1] = work_size[0]; 
			work_size[2] = work_size[0];

			ndrange_t ndrange1 = ndrange_3D(work_size);
			void(^func_SamplingBlock)(void) = ^{ f_RedixSort_OffsetB(Hash, HashIndex, rangeOffset, rangeNum, eNum, Step, offset); };
			int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
			//if(err_ret<0) printf("Previous copy D  = %u \n", err_ret);
		} else {
			work_size[0] = pow(eNum, 1.0f/3.0f)+1; 
			work_size[1] = work_size[0]; 
			work_size[2] = work_size[0]; 
			ndrange_t ndrange1 = ndrange_3D(work_size);
			void(^func_SamplingBlock)(void) = ^{ f_RedixSort_Acc(Hash, HashIndex, rangeOffset, rangeNum, eNum, Step); };
			int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);

		}
		
		
	}
}

void f_RedixSort_OffsetB(	__global ulong2* Hash, __global uint2* HashIndex, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uchar Step, uint offset){
	const size_t id = CurrentID;
	size_t CurrentOffset = id+offset;
	
	if(CurrentOffset<eNum){
		CurrentOffset *= BlockNum;
		for(uchar i=0; i<16; ++i){
			rangeOffset[CurrentOffset+i].s1 = rangeOffset[CurrentOffset+i].s0;
		}
	}
	//barrier(CLK_GLOBAL_MEM_FENCE);
	//===============偏移===================
	if(id==0){
		size_t work_size[3];
		work_size[0] = get_num_groups(0);
		work_size[1] = get_num_groups(1);
		work_size[2] = get_num_groups(2);
		ndrange_t ndrange1 = ndrange_3D(work_size);
		void(^func_SamplingBlock)(void) = ^{ f_RedixSort_OffsetA(Hash, HashIndex, rangeOffset, rangeNum, eNum, Step, offset*2); };
		int err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
		//if(err_ret<0) printf("Previous copy D  = %u \n", err_ret);
	}
}

void f_RedixSort_Acc(__global ulong2* Hash, __global uint2* HashIndex, __global uint2* rangeOffset, __global uchar* rangeNum, uint eNum, uchar Step){
	const size_t id = CurrentID;
	if(id<eNum){
		uint totalOffset = 0;
		uint currentOffset = id*BlockNum;
		uint currentOffset_s1 = id*BlockNum;
		uint indexStep = 0;
		uint endPos = (eNum-1)*BlockNum;
		
		for(uchar i=0; i<16; ++i){
			uchar RangeSize = rangeNum[currentOffset];
			//if(totalOffset >= eNum*BlockNum) printf("err totalOffset  = %u \n", rangeOffset[endPos+i].s0);
			if(RangeSize){
				uint currentStep = rangeOffset[currentOffset].s0 - RangeSize + totalOffset;
				for (uint ia = 0; ia < RangeSize; ++ia) {
					Hash[     currentStep + ia].s0 = Hash[     currentOffset_s1 + indexStep + ia].s1;
                    HashIndex[currentStep + ia].s0 = HashIndex[currentOffset_s1 + indexStep + ia].s1;
				}
				indexStep += RangeSize;
			}
			
			totalOffset += rangeOffset[endPos+i].s0;
			++currentOffset;
		}
	}

	

}





